Skip to content

[Metal] Enable Metal 4 shader compilation#19595

Open
oraluben wants to merge 1 commit into
apache:mainfrom
oraluben:tvm-for-cooperative-tensor
Open

[Metal] Enable Metal 4 shader compilation#19595
oraluben wants to merge 1 commit into
apache:mainfrom
oraluben:tvm-for-cooperative-tensor

Conversation

@oraluben

Copy link
Copy Markdown
Contributor

Use the latest SDK language version when available so generated MPP tensor_ops shaders can compile on M5+ while preserving the existing fallback for older SDKs.

@gemini-code-assist gemini-code-assist Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Code Review

This pull request introduces a preprocessor check to support MTLLanguageVersion4_0 in the Metal runtime. Feedback indicates that since MTLLanguageVersion constants are enums and not macros, the current check will always evaluate to false, rendering the change ineffective. It is suggested to use SDK version macros for proper feature detection.

Comment thread src/runtime/metal/metal_module.mm Outdated
@tlopex

tlopex commented May 25, 2026

Copy link
Copy Markdown
Member

Seems that there is something wrong with CI. Could you retrigger it? @oraluben

@tlopex tlopex left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for updating the feature detection to use the SDK version macro.

I’m concerned that selecting MTLLanguageVersion4_0 solely based on the build SDK can introduce a runtime compatibility regression. A TVM binary built with a macOS 26 SDK may still run on systems/devices where Metal 4 is not supported, but this change would request MSL 4.0 for every Metal runtime compilation, including shaders that would otherwise compile fine with MTLLanguageVersion2_3.

Could we avoid making Metal 4 the unconditional language version whenever the build SDK exposes it? For example, either add a runtime/device capability guard, or try Metal 4 only when needed and fall back to MTLLanguageVersion2_3 if compilation fails.

Also, please clarify what environment was used to validate the new branch, since the current CI does not appear to exercise the __MAC_OS_X_VERSION_MAX_ALLOWED >= 260000 path.

@oraluben oraluben force-pushed the tvm-for-cooperative-tensor branch from 08dd5b4 to 5456e87 Compare June 20, 2026 07:53
@oraluben

oraluben commented Jun 21, 2026

Copy link
Copy Markdown
Contributor Author

Updated the branch to use MSL 4 only when both the build SDK exposes it and the runtime device reports MTLGPUFamilyMetal4; otherwise it keeps the existing MTLLanguageVersion2_3 path.

Validation was done through TileLang PR tile-ai/tilelang#2252. That PR tests the Metal 4 cooperative-tensor path on supported hardware, and also passed TileLang’s Metal tests on GitHub Actions with macOS 26 + M1. The latter is the compatibility case here: the SDK/runtime can expose Metal 4, but M1 does not support MTLGPUFamilyMetal4, so passing those tests confirms the fallback to MSL 2.3 works.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants